{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# PyOpenCL: An exercise"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "import pyopencl as cl\n",
    "import numpy as np\n",
    "import numpy.linalg as la\n",
    "import pyopencl.array\n",
    "import pyopencl.clrandom"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Change the code below to:\n",
    "    \n",
    "* Compute $c_i = a_ib_i$\n",
    "* Use work groups of $16\\times 16$ items\n",
    "* Benchmark $1\\times 1$ workgroups against $16\\times 16$ workgroups\n",
    "\n",
    "  * Use `time()` from the `time` module. (i.e. `import time`)\n",
    "  * Use `queue.finish()`."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "ctx = cl.create_some_context()\n",
    "queue = cl.CommandQueue(ctx)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "a = np.random.rand(1024, 1024).astype(np.float32)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": [
    "prg = cl.Program(ctx, \"\"\"\n",
    "    __kernel void twice(__global float *a)\n",
    "    {\n",
    "      int gid0 = get_global_id(0);\n",
    "      int gid1 = get_global_id(1);\n",
    "      int i = gid1 * 1024 + gid0;\n",
    "      a[i] = 2*a[i];\n",
    "    }\n",
    "    \"\"\").build()\n",
    "twice = prg.twice"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {
    "collapsed": false
   },
   "outputs": [
    {
     "data": {
      "text/plain": [
       "<pyopencl.cffi_cl.Event at 0x7f679a3145f8>"
      ]
     },
     "execution_count": 5,
     "metadata": {},
     "output_type": "execute_result"
    }
   ],
   "source": [
    "a_dev = cl.array.to_device(queue, a)\n",
    "twice(queue, a_dev.shape, None, a_dev.data)"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {
    "collapsed": false
   },
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "0.0 591.347\n"
     ]
    }
   ],
   "source": [
    "print(la.norm(a_dev.get() - 2*a), la.norm(a))"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": false
   },
   "outputs": [],
   "source": []
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.5.0+"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 0
}
